__device__ int addVar(int* rowStore, int addingLamVar, int maxSize, int* rep) {
	for(int i = 1; i < rowStore[0]; i++) {
		if(rowStore[i] == addingLamVar) {
			return 0;
		}
	}
	unsigned int tail;
	
	atomicAdd(&rep[0],1);
	tail = atomicAdd(&rowStore[0],1);
	if(rowStore[0] > maxSize) {
		atomicAdd(&rowStore[0],-1);
		return 1;
	} else {
		rowStore[tail] = addingLamVar;
	}

	return 0;
}


extern "C"  __global__ void eigenCFA(int* Fun, int* Arg1, int* Arg2, int* Store, int* semaphore, int* rep, int scaleM, int scaleCall, int scaleLam) {
	int column = blockDim.x*blockIdx.x + threadIdx.x;

	if(column < scaleCall && blockIdx.y < 2) {//y?

		int Argi[2];
		int numCall = column;
		int* L = &Store[Fun[numCall]*scaleM];
		
		Argi[0] = Arg1[numCall];
		Argi[1] = Arg2[numCall];

		int* Li = &Store[Argi[blockIdx.y]*scaleM];

		int vL[2];

		int error = 0;

		int numV = numCall*scaleM;
		for(int j = 1; j < L[0]; j++) {
			if(L[j] < 0)
			{
				break;
			}
			vL[0] = L[j] - 1;
			vL[1] = L[j] - 1 + scaleLam;
			for(int k = 1; k < Li[0]; k++) {
				if(Li[k] < 0) {
					break;		
				}
				bool leaveLoop = false;
				while(!leaveLoop) {
					if(atomicExch(&(semaphore[vL[blockIdx.y]]), 1u) == 0u) {
						error = addVar(&Store[vL[blockIdx.y]*scaleM], Li[k], scaleM, rep);
						
						leaveLoop = true;
						atomicExch(&(semaphore[vL[blockIdx.y]]), 0u);
		//				if(error) {
		//					printf("memEnd vL[%d]!\n", blockIdx.y);
		//				} 
					}
				}
			}
		}
	}
}

/*__device__ int addVar(int* rowStore, int addingLamVar, int maxSize, int* rep) {
	int numCall = rowStore[0];
	for(int i = 1; i < rowStore[0]; i++) {
		if(rowStore[i] == addingLamVar) {
			return 1;
		}
	}

	unsigned int tail;
	//rep[0] = 1;
	atomicAdd(&rep[0], 1);
	//rowStore[numCall] = addingLamVar;
	//rowStore[0]++;
	tail = atomicAdd(&rowStore[0], 1);
	rowStore[tail] = addingLamVar;
	return 2;
}


extern "C"  __global__ void eigenCFA(int* Fun, int* Arg1, int* Arg2, int* Store, int* semaphore, int* rep, int scaleM, int scaleCall, int scaleLam) {

	int column = blockDim.x*blockIdx.x + threadIdx.x;

	if(column < scaleCall && blockIdx.y < 2) {

		int Argi[2];

		int numCall = column;
		int* L = &Store[Fun[numCall]*scaleM];
		
		Argi[0] = Arg1[numCall];
		Argi[1] = Arg2[numCall];

		int* Li = &Store[Argi[blockIdx.y]*scaleM];
		int vL[2];
		int error = 0;

		//int startL = L[0];
//		int j = 1;
//		int k = 1;
		//while(L[j] > 0) {
		int start = L[0];
		for(int j = 1; j < start; j++) {
			if(L[j] < 0) {
				break;
			}
			vL[0] = L[j] - 1;
			vL[1] = vL[0] + scaleLam;
			int startLi = Li[0];
			for(int k = 1; k < startLi; k++) {
				if(Li[k] < 0) {
					break;
				}
	//			while(Li[k] > 0) {
//				bool leaveLoop = false;
//				while (!leaveLoop) {
//					if (atomicExch(&(semaphore[vL[blockIdx.y]]), 1u) == 0u) {
						error = addVar(&Store[vL[blockIdx.y]*scaleM], Li[k], scaleM, rep);
						//critical section
	//						leaveLoop = true;
//						atomicExch(&(semaphore[vL[blockIdx.y]]),0u);
					}
				} 
			}
		}
	}
}
*/